__kernel void scan(__global int * g_in,
				   __global int * g_out,
				   __local int *l_temp,
				   int n) {

	int thid = get_local_id(0);  
	int offset = 1;  

	l_temp[2*thid] = g_in[2*thid]; // load input into shared memory  
	l_temp[2*thid+1] = g_in[2*thid+1];  
 	
	for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree  
	{   
		barrier(CLK_LOCAL_MEM_FENCE); 
		if (thid < d)  
		{  
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
			l_temp[j] += l_temp[i];  
		}  
		offset *= 2;
	}	
    
	// clear the last element  
	if (thid == 0) 
	{ 
		l_temp[n - 1] = 0; 
	} 
                 
	for (int d = 1; d < n; d *= 2) // traverse down tree & build scan  
	{  
		offset >>= 1;  
		barrier(CLK_LOCAL_MEM_FENCE);  
		if (thid < d)                       
		{   
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
 	
			float t = l_temp[j];  
			l_temp[j] += l_temp[i];  
			l_temp[i] = t;   
		}  
	}  
	barrier(CLK_LOCAL_MEM_FENCE);  

	g_out[2*thid] = l_temp[2*thid]; // write results to device memory  
	g_out[2*thid+1] = l_temp[2*thid+1];  
}